#include "SharedMem.h"
#include <cuda_runtime_api.h>
#include <device_functions.h>
#include "device_launch_parameters.h"
#include <iostream>
#include <time.h>
#include <cuda.h>
using namespace std;

namespace conflict_test
{
const int SHARED_ARR_SIZE = 1024;
const int ROUNDS = 1000;
}

__global__ void complete_conflict(double* output){
	__shared__ double arr [conflict_test::SHARED_ARR_SIZE];

	arr[threadIdx.x] = blockIdx.x;

	/*if(threadIdx.x == 0){
		for(int i = 0; i < conflict_test::SHARED_ARR_SIZE; i++) arr[i] = blockIdx.x;
	}*/

	__syncthreads();

	double result = 0;
	for(int i = 0; i < conflict_test::ROUNDS; i++){
		for(int j = 0; j < conflict_test::SHARED_ARR_SIZE; j++){
			result += arr[0];
		}
	}   

	result /= (conflict_test::ROUNDS * conflict_test::SHARED_ARR_SIZE);

	if(threadIdx.x == 0){
		output[blockIdx.x] = result;
	}
}

__global__ void avoid_conflict(double* output){
	__shared__ double arr [conflict_test::SHARED_ARR_SIZE];

	arr[threadIdx.x] = blockIdx.x;

	/*if(threadIdx.x == 0){
		for(int i = 0; i < conflict_test::SHARED_ARR_SIZE; i++) arr[i] = blockIdx.x;
	}*/

	__syncthreads();

	int offset = 1*threadIdx.x;

	double result = 0;
	for(int i = 0; i < conflict_test::ROUNDS; i++){
		for(int j = 0; j < conflict_test::SHARED_ARR_SIZE; j++){
			//result += arr[(j+offset)% conflict_test::SHARED_ARR_SIZE];
			result += arr[threadIdx.x];
		}
	}

	result /= (conflict_test::ROUNDS * conflict_test::SHARED_ARR_SIZE);

	if(threadIdx.x == 0){
		output[blockIdx.x] = result;
	}
}

void bank_conflict_test(){
	int num_of_blocks = 100;
	int rounds = 10;

	double* result_device;
	double* result_host = (double*)malloc(num_of_blocks*sizeof(double));
	cudaMalloc((void**)&result_device,num_of_blocks*sizeof(double));

	clock_t start = clock();
	for(int i = 0; i < rounds; i++){
		complete_conflict<<<num_of_blocks,128>>>(result_device);
	}

	cudaMemcpy(result_host,result_device,num_of_blocks*sizeof(double),cudaMemcpyDeviceToHost);
	cout << "conflict version takes: " << ((double)clock()-start)/CLOCKS_PER_SEC << endl;
	cout << "result 1: " << endl;
	for(int i = 0; i < num_of_blocks; i++) cout << result_host[i] << " ";
	
	cout << endl << endl;;

	start = clock();
	for(int i = 0; i < rounds; i++){
		avoid_conflict<<<num_of_blocks,128>>>(result_device);
	}

	cudaMemcpy(result_host,result_device,num_of_blocks*sizeof(double),cudaMemcpyDeviceToHost);
	cout << "conflict version takes: " << ((double)clock()-start)/CLOCKS_PER_SEC << endl;
	cout << "result 2: " << endl;
	for(int i = 0; i < num_of_blocks; i++) cout << result_host[i] << " ";

	cudaFree(result_device);
	free(result_host);
}